fix(KERNEL-SSM-MAMBA): the derived bound omitted nvcc FMA contraction, and the memory-safety claim outran the kernels (#496) by localai-bot · Pull Request #592 · mudler/vllm.cpp

localai-bot · 2026-08-13T07:00:39Z

Tightening pass on the fresh review of #566 (row/KERNEL-SSM-MAMBA-SSD-W2-FINISH @ 1e819144e, W2 of .agents/specs/mamba2-ssd.md). That review's verdict was PASS; these are its two MEDIUM and three LOW findings. Branched from 1e819144e with origin/main merged, so it is standalone and does not stack on #566.

Issue: #496.

The review verified the equivalence contract's central MMA reasoning, reconstructed the reformulated mutations on CPU twins and proved each reds, and confirmed the grep, the recovered-byte md5s and the CI baseline subtraction. None of that is re-done here.

F1 (MEDIUM) — the derived bound's derivation omitted nvcc FMA contraction

The header and both test headers claimed the elementary functions were the only admitted source of device-vs-host divergence. False as written. Host C++ is pinned -ffp-contract=off (CMakeLists.txt:55) so a*b + c keeps two roundings; nothing passes --fmad=false to nvcc, so cuda_mamba2_ssd.cuh compiles at the default --fmad=true and every acc += a*b in it is a single-rounding fma whose host twin is not. This repo has already measured that idiom (.agents/benchmark-record.md:532). CMakeLists.txt:41-56 carves CUDA out of the contraction policy because "GPU parity tests compare GPU-vs-GPU" — G2 is exactly the case that carve-out does not cover.

The arithmetic:

term	bound	source
libm	`2.5·K·u`	≤ 2.5 ulp per decay factor (CUDA `expf` ≤ 2, glibc ≤ 0.5), through a product of at most `K`
summation	`(K-1)·u`	length-`K` f32 sum; this is what amplifies the libm term
contraction	`K·u`	the `K` product roundings the host keeps and the device's `fma` does not
total	`4.5·K·u − u`

old 4·(K+2)·u: 4.5K − 1 ≤ 4K + 8 ⟺ K ≤ 18. Not provable at the driver shapes, which run at K = T = 200.
new 5·(K+2)·u: 4.5K − 1 ≤ 5K + 10 ⟺ 0.5K + 11 ≥ 0. All K ≥ 0.

-fmad=false was weighed and REJECTED, not overlooked: nvcc takes it per translation unit and this is a header included by cuda_gdn.cu:48, so applying it means de-contracting every GDN decode kernel in that TU (a measured hot path) or splitting a new src/vt/ TU, which §8.3 already records as blocked on #515. Slowing a shipped kernel to make a bound's prose true is the wrong trade.

Nothing was hidden numerically. Re-scaling §8.4's audit by 4/5: worst of 55 comparisons 7.66% → 6.13% of budget, driver shapes 0.32%/0.18% → 0.26%/0.14%, mutant M3 962173% → 769738% — still caught by four orders of magnitude. The constant moved because the derivation gained a term the build emits. The header comment, all three test comments and spec §8.3 now say the same true thing.

F2 (MEDIUM) — the memory-safety claim outran the kernels

Both halves of the finding are taken: the clamps AND the narrowed claim.

The shared validator checks metadata shape/dtype/device only (CheckI32Meta, ops.cpp:1717-1723); every value check lives in the CPU kernel (cpu_ops.cpp:1622-1648), so the device arm silently drops six. The stated reason for dropping them — a D2H plus a stream sync — does not apply to the memory-unsafe ones: those values are already in device registers, and the decode kernel has always clamped its state_indices slot for free on that basis.

Clamped:

M2StatePassKernel — chunk_end to nchunks, chunk_start to 0. Unclamped, lci[b] >= nchunks makes M2Store(passed, …) an out-of-bounds write past the cudaMallocAsync allocation (W1 finding F7's device half).
M2ChunkScanKernel — si_ok = si >= 0 && si < S, and !si_ok opens the chunk from a zero previous state. Unclamped, an out-of-range seq_idx[c] reads initial_states out of bounds, and a seq_idx[0] < 0 makes si == si_prev at c == 0 and indexes passed at chunk −1 — a hole the finding did not name, found while writing the clamp. The kernel gained an S parameter for it.

In-contract behaviour is bit-identical: in contract si is always in range and lci always below nchunks, so neither clamp can fire.

Narrowed: clamping two does not make the arm memory-safe. The cu_chunk_seqlens tiling and per-chunk length checks are NOT clamped and a violation IS memory-unsafe (a garbage ccs indexes x/B/C/z/out out of bounds in every stage). The header and §8.3 now say that, instead of a blanket "the device kernels remain MEMORY SAFE".

Pinned, not asserted in prose: a device-only case runs both violations against in-contract reference runs whose result each clamp is defined to reproduce, so the assertions are exact rather than tolerances.

F3 / F4 (LOW, record accuracy)

M7's §8.4 label overstated the device mutant: the launcher's nblocks = rows * args.n_groups is not mutated while the kernel's n_groups is forced to 1, so blocks blk >= rows run past the tensor. The mutant is memory-unsafe and fails partly for that, not purely on whole-row variance. The guarantee is pinned by the reviewer's clean CPU twin; only the label was wrong.
§8.2's residual still gave the downcasts as the reason W2 cannot byte-compare, written when W2 was expected to mirror them. §8.3 supersedes it: both arms stay f32 and the reasons are libm and contraction. Two sentences, two causes, one fact — reconciled.

F5 (LOW) — taken

The five per-call scratch buffers are held by an M2Scratch scope guard, so a throw on the Nth cudaMallocAsync no longer leaks the N−1 before it; Release() also frees all five before reporting rather than leaking the remainder on a mid-sequence free failure.

Evidence (implementer, CPU box)

df -h / logged beside every result: 83–87% used / 60–73G free throughout.

suite	test cases	assertions	status
`test_ops_mamba2_ssd`	8 / 8 passed	1175 / 1175	`SUCCESS!`
`test_ops_mamba2_state_update`	6 / 6 passed	2469 / 2469	`SUCCESS!`
`test_ops_mamba2_gated_norm`	9 / 9 passed	2107 / 2107	`SUCCESS!`

Identical to the pre-change counts, as expected — every code change is inside #ifdef VLLM_CPP_CUDA or in the .cuh. Status: was read, not assertions: alone. Clean Release rebuild: 817/817 edges, 0 warnings, 0 errors under -Wall -Wextra -Werror. Full ctest: 403/403 passed, 0 failed. scripts/agent-preflight.sh --staged: RC=0.

The CUDA arm COMPILES under real nvcc — operator-run, not CI

dgx.casa returned at 06:57 UTC after a reboot. The operator (not me — this box has no nvcc and no GPU) transferred this branch by git archive (never rsync) and built it with real nvcc:

NVCC: cuda_13.0.r13.0
CUTLASS found at ~/cutlass-4.5.0; enabling sm120a NVFP4 cutlass GEMM
Marlin NVFP4 W4A16 MoE GEMM enabled (vendored) for [121a]
FlashAttention-2 prefill/decode: ENABLED for arch(es) [121a]
CONFIGURE_EXIT=0   BUILD_EXIT=0   WARNINGS=0   ENOSPC=0

Zero errors, zero warnings under -Werror, disk unchanged at 64G either side (so not the stale-binary false green), and the three fast-path features read out of the configure log rather than assumed — an absent CUTLASS also exits 0, so "the build succeeded" alone would prove nothing.

That retires the risk the shim check below could only approximate: the new S kernel parameter and the M2Scratch guard compile through cuda_gdn.cu at the shipping arch. Attributed to the operator, not to CI — cuda-fat-build has still never completed on this branch.

⚠️ The red checks on this PR are CANCELLED, not FAILED

Every Actions run here — four SHAs, both workflows — ended conclusion: cancelled. Three were my own follow-up pushes. The fourth was killed at 07:44:45–07:44:55 together with every run in the repository (20 of 20 across 7 branches); windows-msvc-cpu died mid-build after passing two steps. That is an Actions-side event, not a verdict on this diff.

gh pr checks renders a cancelled job as fail. My watcher reported 16 failures and the operator's reported 20; the true count both times was zero. Only the run-level conclusion distinguishes the two — gh run view <id> --json status,conclusion. Whatever the current run reports is the CI result; nothing here claims one.

`omitted_gates` — blocked by the GPU LOCK, not the host

Two substitutes were run by the implementer before the operator's compile existed; neither is offered as the device gate:

.cuh compile + arity check. Compiled at -std=c++20 -Wall -Wextra -Werror against CUDA shims, with each Kernel<<<cfg>>>(args) rewritten to M2Sink(cfg), Kernel(args) — dropping the launch config while preserving the argument-count and type check on all 7 launches. EXIT 0. Proved armed by deleting the S argument the F2 clamp added, on a scratch copy: too few arguments to function M2ChunkScanKernel, exit 1. Superseded by the operator's real nvcc build above.
F2 clamp CPU twin. The two index computations walked over the device case's own shape. Unclamped, every claimed hole reproduced: index 2432 and −4096 into a 2048-element passed, 463232 into a 2048-element initial_states, −512 for the c == 0 hole. Clamped, all land in [0, 1920], the lci clamp reproduces the in-contract index range exactly, and both seq_idx violations read no previous state at all. Not superseded by any compile.

Still owed — and the blocker is now $HOME/gpu.lock, not the machine. The host is back; the lock is held by other coordinators' jobs with 8-hour timeouts, and running against a contended GPU would reproduce the undetected-contention defect §8.4 already records.

the three suites' CUDA arms (Release and Debug) — execution, which a compile does not supply;
compute-sanitizer memcheck on the new clamp case. This is what actually proves F2: an out-of-bounds write into a cudaMallocAsync pool commonly does not fault, so a green run is necessary and not sufficient;
the 9-mutation re-sweep against the moved bound 5·(K+2)·u — §8.4's sweep was scored against 4·(K+2)·u, and a widened bound is precisely the change that could stop a mutation reddening. The re-scaled margins make it very likely to hold; very likely is not a result;
§8.4's ~/w2ssd/refail.log, still REMOTE_UNVERIFIED.

The two Windows CI jobs are the main baseline established in §8.4 (#512 fixed by #583, #514, #584) — subtracted, not inherited.

🤖 Generated with Claude Code

…ps (#496) W2 of .agents/specs/mamba2-ssd.md. These are the failing tests, committed before the kernels they gate, per the implementer contract. The three suites gain a `#ifdef VLLM_CPP_CUDA` section that runs the SAME inputs through the device arm. They fail for the intended reason: no NATIVE kernel is registered for kMamba2ChunkScan / kMamba2StateUpdate / kRmsNormGatedGroup on DeviceType::kCUDA. That reason is NOT "GetOp throws", and the difference is the whole point of one assertion in these suites. GB10 is `integrated && pageable_memory_access` (cuda_backend.cu Registrar), so `Backend::UnifiedMemory()` is TRUE and `ReferenceTierEligible(kCUDA)` with it. On a GetOp miss the provider seam does not throw: it installs the CPU HOST kernel as a `kReferenceProviderName` provider and runs THAT over the device pointers (op_provider.h, "portable reference tier"). Every numeric assertion in a device arm would then pass while nothing ran on the GPU -- the device arm gated by running the host arm twice. So every CUDA case calls `RequireNativeCudaProvider`, which reads `GetOpProviderStats(op, kCUDA).last_selected` and refuses `vt-cpu-ref`. These are EAGER dispatches, not a captured graph, so the counter is genuinely populated ([[graph-replay-does-no-host-dispatch-counters-read-zero]]). The declared equivalence contract is written down here BEFORE the kernel, in the head comment of the SSD suite's CUDA section: * the CUDA arm keeps f32 accumulation throughout and does NOT mirror the tile downcasts in upstream's Triton dots (ssd_chunk_state.py:283-285, ssd_chunk_scan.py:266-269, :359-363) -- those are the input-precision requirement of `tl.dot`, i.e. of a tensor-core MMA, and every one of those tiles is loaded `.to(tl.float32)` and computed in f32 right up to the MMA. The memory format is unchanged, so this is not a "too wide" dtype; * G1, the primary gate, is the device output against the SAME independent double-precision sequential reference at the SAME upstream-ported tolerances the host arm is held to; * G2, device-vs-host, is a DERIVED bound: `rtol(K) = 4*(K + 2)*2^-24` over a recurrence of length K. CUDA's `expf` is documented to <= 2 ulp and glibc's to <= 0.5, so a product of K decay factors carries <= 2.5*K*u of libm disagreement, and the length-K f32 summation adds the standard (K-1)*u -- 3.5*K*u, rounded up to integers. Everything else is held identical by construction: each device output element is accumulated in ONE thread over the host arm's index range in the host arm's direction, so summation order is not a second source. A BYTE COMPARE IS NOT REACHABLE, and the libm difference is exactly why. The slack actually used is REPORTED on every comparison, so a bar that stopped doing work would be visible rather than silently absorbing a defect. Also lands the mutation-proof §8.2 records as owed: the decode kernel's `CheckMamba2ANegative` at cpu_ops.cpp:1877 was pinned by NO test -- deleting it left test_ops_mamba2_state_update fully green while the same deletion on its chunk-scan twin reds. The new "A must be negative" SUBCASE mirrors test_ops_mamba2_ssd.cpp:900 and additionally pins that the guard is a SIGN test, not an accidental magnitude floor (A = -1e-30 is accepted). FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [Claude Code]

…he gate host after a worktree loss (#496) FOLLOWING_AGENTS_PROTOCOL NOT AUTHORED BY THE COMMITTER. This is an operator recovery of a fresh implementer's work after the fourth external deletion of an isolation worktree this session. The implementer had built and run this green on both boxes and staged it; the worktree was removed before the commit. The bytes survived on the gate host and are restored here unchanged: md5 cbb1f928f4 for cuda_mamba2_ssd.cuh and 37a0404433 for cuda_gdn.cu, matching what the implementer reported before the loss. The declared equivalence contract, which the implementer decided BEFORE writing the kernel and recorded in the kernel header and all three test headers: The CUDA arm keeps f32 accumulation throughout and deliberately does NOT mirror upstream's tile downcasts. Those casts -- b.to(x_ptr.dtype.element_ty) at ssd_chunk_state.py:283-285, cb.to(...)/prev_states.to(...) at ssd_chunk_scan.py:266-269,359-363 -- are the input-precision requirement of tl.dot, a tensor-core MMA. Every one of those tiles is loaded .to(tl.float32) and computed in f32 right up to the MMA. These are scalar-FMA kernels with no MMA, so mirroring the downcast would copy a constraint we do not have. The inter-chunk `passed` buffer is allocated at state_dtype, NOT the host arm's f32 working width that spec 8.2 F9 warned W2 must not inherit. A byte compare against the host arm is NOT reachable, and the downcasts are not why: the two arms call different libms (CUDA expf <= 2 ulp, glibc <= 0.5). Everything else is identical by construction. So the primary gate is the device output against the same double-precision sequential reference at the same upstream-ported tolerances the host arm uses, on the same inputs -- which separates "device defect" from "wrong threshold". The derived device-vs-host bar is rtol(K) = 4*(K+2)*2^-24, derived from 2.5 ulp of libm disagreement per decay factor through a product of at most K plus (K-1)*u summation error. No number was tuned and no tolerance was widened; each comparison logs the fraction of budget actually used through MESSAGE rather than INFO, because doctest prints INFO only on failure and an unaudited bar would have been a false claim. Evidence already captured on the gate host: Release build for 121a with CUTLASS 4.5.0, fa2 ENABLED and Marlin NVFP4 enabled, 0 warnings; RED run SIGSEGV on all three binaries; GREEN run ssd 11 cases / 2069 assertions, state_update 10 / 5965, gated_norm 12 / 3723, all Status SUCCESS, exit 0, with zero reference-tier lines. That RED SIGSEGV is a real shared-seam defect, filed as #547 and deliberately not fixed in flow: GB10 reports Backend::UnifiedMemory() == true, so ReferenceTierEligible(kCUDA) is true, and with no native kernel GetOp installs the CPU host kernel as a vt-cpu-ref provider and runs it over cudaMalloc pointers. include/vt/backend.h already says a cudaMalloc pointer is not host-dereferenceable on GB10; op_provider.cpp:515-526 gates on UnifiedMemory() where it needs DeviceMemoryIsHostAddressable(). Every CUDA case here now calls RequireNativeCudaProvider, so a device arm can never be gated by running the host arm twice. STILL OWED, and this branch is NOT landable until a fresh implementer finishes it: the 8 scripted CUDA mutations, compute-sanitizer, the Debug arm, a full ctest on the gate host, the spec section 8.3 that records the contract above (its only copy was a staged blob and is presumed lost), and an origin/main re-merge. A fresh review follows that, not this commit. Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [Claude Code]

FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [Claude Code]

…e $8.2 decode SUBCASE (#496) FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [Claude Code]

FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [Claude Code]

…aught, sanitizer clean, Debug arm green (#496) FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [Claude Code]

…tributed failure (#496) FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [Claude Code]

… -- a second job locks /tmp/gpu.lock (#496) FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [Claude Code]

…e attribution re-run is REMOTE_UNVERIFIED (#496) FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [Claude Code]

FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [Claude Code]

…, and the memory-safety claim outran the kernels (#496) Tightening pass on PR #566 (`row/KERNEL-SSM-MAMBA-SSD-W2-FINISH` @ `1e819144e`), whose fresh review returned PASS with five findings. The review's own work -- the equivalence contract's MMA reasoning, the reformulated mutations on CPU twins, the grep, the recovered-byte md5s, the CI baseline subtraction -- is not re-done here. F1 (MEDIUM) -- the derivation, not the number, was wrong. `cuda_mamba2_ssd.cuh` and the test headers both claimed the elementary functions were the ONLY admitted source of device-vs-host divergence. They are not. Host C++ is pinned `-ffp-contract=off` (CMakeLists.txt:55) so `a*b + c` keeps two roundings; nothing passes `--fmad=false` to nvcc, so the header compiles at the default `--fmad=true` and every `acc += a*b` in it is a single-rounding `fma` whose host twin is not. This repo has MEASURED that idiom already (.agents/benchmark-record.md:532). CMakeLists.txt:41-56 carves CUDA out of the contraction policy because "GPU parity tests compare GPU-vs-GPU"; G2 is exactly the case that carve-out does not cover. The arithmetic, which the old constant did not survive: libm 2.5*K*u (<= 2.5 ulp per decay factor, CUDA expf <= 2, glibc <= 0.5) summation (K-1)*u (length-K f32 sum; this is what amplifies the libm term) contraction K*u (the K product roundings the host keeps and fma does not) ---------------------------------------------------------------------------- total 4.5*K*u - u old 4*(K+2)*u: 4.5K - 1 <= 4K + 8 <=> K <= 18. NOT PROVABLE at the driver shapes, which run at K = T = 200. new 5*(K+2)*u: 4.5K - 1 <= 5K + 10 <=> 0.5K + 11 >= 0. All K >= 0. `-fmad=false` was weighed and REJECTED, not overlooked: nvcc takes it per translation unit and this is a header included by `cuda_gdn.cu:48`, so applying it means de-contracting every GDN decode kernel in that TU -- a measured hot path -- or splitting a new `src/vt/` TU, which §8.3 already records as blocked on #515. Slowing a shipped kernel to make a bound's prose true is the wrong trade. Nothing was hidden numerically: re-scaling §8.4's audit by 4/5, the worst of 55 comparisons goes 7.66% -> 6.13% of budget, the driver shapes 0.32%/0.18% -> 0.26%/0.14%, and mutant M3 962173% -> 769738%, still caught by four orders of magnitude. The bound moved because the DERIVATION gained a term the build actually emits, and the header comment, all three test comments and spec §8.3 now say the same true thing. F2 (MEDIUM) -- both halves taken: the free clamps AND the narrowed claim. The shared validator checks metadata shape/dtype/device only (`CheckI32Meta`, ops.cpp:1717-1723); every VALUE check lives in the CPU kernel (cpu_ops.cpp:1622- 1648), so the device arm silently drops six of them. Three are memory-unsafe, and the stated reason for dropping them -- a D2H plus a stream sync -- does not apply, because the values are already in device registers and the decode kernel has always clamped its `state_indices` slot for free on that basis. * `M2StatePassKernel`: `chunk_end` clamped to `nchunks`, `chunk_start` to 0. Unclamped, `lci[b] >= nchunks` makes `M2Store(passed, ...)` an out-of-bounds WRITE past the `cudaMallocAsync` allocation (W1 finding F7's device half). * `M2ChunkScanKernel`: `si_ok = si >= 0 && si < S`, and `!si_ok` opens the chunk from a zero previous state. Unclamped, an out-of-range `seq_idx[c]` reads `initial_states` out of bounds, and a `seq_idx[0] < 0` additionally makes `si == si_prev` at c == 0 and indexes `passed` at chunk -1 -- a hole the finding did not name and this pass found while writing the clamp. The kernel gained an `S` parameter for it. In-contract behaviour is bit-identical: in contract `si` is always in range and `lci` always below `nchunks`, so neither clamp can fire. The claim is narrowed at the same time, because clamping two does not make the arm memory-safe. The `cu_chunk_seqlens` tiling and per-chunk length checks are NOT clamped and a violation IS memory-unsafe -- a garbage `ccs` indexes x/B/C/z/out out of bounds in every stage -- and the header and §8.3 now say so instead of folding it into a blanket "the device kernels remain MEMORY SAFE". Pinned, not asserted in prose: a device-only case runs both violations against in-contract reference runs whose result each clamp is DEFINED to reproduce, so the assertions are exact rather than tolerances. F3, F4 (LOW, record accuracy). M7's §8.4 label overstated the device mutant: the launcher's `nblocks = rows * args.n_groups` is not mutated while the kernel's `n_groups` is forced to 1, so blocks `blk >= rows` run past the tensor and the mutant is memory-unsafe, failing partly for that rather than purely on whole-row variance. The guarantee IS pinned by the reviewer's clean CPU twin; only the label was wrong. §8.2's residual sentence still gave the DOWNCASTS as the reason W2 cannot byte-compare, written when W2 was expected to mirror them; §8.3 supersedes it -- both arms stay f32 and the reasons are libm and contraction. Two sentences in one spec gave two causes for one fact; reconciled. F5 (LOW). Taken. The five per-call scratch buffers are held by an `M2Scratch` scope guard, so a throw on the Nth `cudaMallocAsync` no longer leaks the N-1 before it, and `Release()` frees all five before reporting rather than leaking the remainder on a mid-sequence free failure. EVIDENCE. `df -h /` 85% used / 67G free before and after every result. test_ops_mamba2_ssd 8/8 1175/1175 SUCCESS! test_ops_mamba2_state_update 6/6 2469/2469 SUCCESS! test_ops_mamba2_gated_norm 9/9 2107/2107 SUCCESS! Identical to the pre-change counts, as expected -- every code change is inside `#ifdef VLLM_CPP_CUDA` or in the `.cuh`. `Status:` was read, not `assertions:` alone. OMITTED_GATES -- the CUDA arm was neither built nor run. `dgx.casa` has been unreachable since 06:50 CEST and this box has no nvcc and no GPU. Two substitutes were run and neither is offered as the device gate: 1. The `.cuh` compiled at `-std=c++20 -Wall -Wextra -Werror` against CUDA shims with each `Kernel<<<cfg>>>(args)` rewritten to `M2Sink(cfg), Kernel(args)` -- dropping the launch config while PRESERVING the arity and type check on all 7 launches. EXIT 0. Proved ARMED by deleting the `S` argument on a scratch copy: `too few arguments to function M2ChunkScanKernel`, exit 1. 2. A CPU twin of the two clamped index computations, over the device case's own shape. Unclamped, every claimed hole reproduced: index 2432 and -4096 into a 2048-element `passed`, 463232 into a 2048-element `initial_states`, -512 for the c == 0 hole. Clamped, all land in [0,1920], the `lci` clamp reproduces the in-contract index range exactly, and both `seq_idx` violations read no previous state at all. Still owed on device: the three CUDA arms, `compute-sanitizer memcheck` on the new case (which is what actually proves memory safety -- an out-of-bounds write into a pool allocation commonly does not fault), a mutation re-sweep against the moved bound, and §8.4's `~/w2ssd/refail.log`, still REMOTE_UNVERIFIED. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [Claude Code]

…g pass (#496) Clean Release rebuild of every target -- 817/817 ninja edges, 0 warnings, 0 errors at `-Wall -Wextra -Werror` -- then full `ctest -j 4`: 100% tests passed, 0 failed out of 403, CTEST_EXIT=0, 22.35 s, 2 skipped. `df -h /` 87% used / 60G free. Recorded with the caveat that matters: this is the CPU-ONLY lane, so it is a much smaller and faster gate than §8.4's 431-test GPU-host run and is not a substitute for it. None of §8.4's ten failures is reachable from a build with no CUDA. The device arm remains `omitted_gates` while `dgx.casa` is unreachable. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [Claude Code]

…D bound (#496) §8.4 reports "the worst one used 7.66% of rtol(K) = 4*(K+2)*2^-24". That is still what the run produced, but §8.5 moved the bound to 5*(K+2)*2^-24, so read without a pointer it now looks like a statement about the current bar. Cross-referenced rather than restated: the captured numbers stay as captured, with the conversion (6.13%, 0.26%/0.14%, 769738%) named next to them and the arithmetic in §8.3 point 6. Rewriting a measurement to match a later derivation is how a record stops being evidence. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [Claude Code]

…ed the gap it does not cover (#496) Two inaccuracies in the comment written one commit ago, found re-reading it: It said "two of the dropped ones are memory-UNSAFE", which reads as "and the other four are not". The header enumerates six, of which the `cu_chunk_seqlens` tiling and per-chunk length checks are ALSO memory-unsafe and are NOT clamped. The two this case covers are the two that are memory-unsafe AND bounded by a register-local clamp. Narrowing the count in the very comment that exists to stop a claim outrunning the kernels would have re-introduced F2 at a smaller scale, so the comment now points at the header's full list and names the uncovered gap explicitly. It also cited §8.4 for the owed `compute-sanitizer memcheck`; that is recorded in §8.5. Comment-only. Rebuilt clean (392/392, 0 warnings) and re-ran: 8/8, 1175/1175, `Status: SUCCESS!`. The `#ifdef VLLM_CPP_CUDA` region was re-checked with a `-DVLLM_CPP_CUDA -fsyntax-only` compile at -Wall -Wextra -Werror, exit 0. `df -h /` 83% used / 73G free. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [Claude Code]

…erator-run on dgx, not CI (#496) Two corrections to §8.5, both narrowing what this branch claims. 1. THE CUDA COMPILE AXIS IS CLOSED, AND IT WAS NOT MY RESULT. `dgx.casa` returned at 06:57 UTC after a reboot, which supersedes §8.4's REMOTE_UNVERIFIED note on the host being away. The operator transferred this branch by `git archive` -- never rsync, which has previously overwritten goldens into false passes -- and built it with real nvcc: NVCC: cuda_13.0.r13.0 CUTLASS found at ~/cutlass-4.5.0; enabling sm120a NVFP4 cutlass GEMM Marlin NVFP4 W4A16 MoE GEMM enabled (vendored) for [121a] FlashAttention-2 prefill/decode: ENABLED for arch(es) [121a] CONFIGURE_EXIT=0 BUILD_EXIT=0 WARNINGS=0 ENOSPC=0 Zero errors, zero warnings under the project's -Werror flags, disk unchanged at 64G either side so this is not the stale-binary false green, and the three fast-path features READ OUT OF the configure log rather than assumed -- an absent CUTLASS exits 0 too, so "the build succeeded" alone proves nothing. That retires the risk my shim arity-check could only approximate: the new `S` kernel parameter and the `M2Scratch` guard compile through `cuda_gdn.cu` at the arch this ships on. Recorded as OPERATOR-RUN. I did not run it, this box has no nvcc, and `cuda-fat-build` has still never completed on this branch -- so no CI job has compiled this code either. The shim check and the F2 clamp CPU twin are kept as what an implementer could establish unaided, not restated as the gate. 2. THE REMAINING GATES ARE BLOCKED BY THE GPU LOCK, NOT BY THE HOST. The old wording said "omitted_gates until dgx.casa returns". It has returned, and they are still owed -- `$HOME/gpu.lock` is held by other coordinators' jobs with 8h timeouts. A reachable host is not an available GPU, and running these against a contended one reproduces exactly the undetected-contention defect §8.4 already records. Owed: the three CUDA arms (EXECUTION, which a compile does not supply); `compute-sanitizer memcheck` on the new clamp case, which is what actually proves F2 because an out-of-bounds write into a cudaMallocAsync pool commonly does not fault; the 9-mutation re-sweep against the MOVED bound, since a widened bound is precisely the change that could stop a mutation reddening; and §8.4's refail.log. Item 3 is the one a reader is most likely to wave through. The re-scaled margins in §8.3 point 6 make it very likely to hold, and very likely is not a result. 3. THE CANCELLED-vs-FAILED TRAP, WHICH TWO OF US HIT FROM OPPOSITE ENDS. Every Actions run for #592 -- four SHAs, both workflows -- ended `conclusion: cancelled`. Three were my own follow-up pushes. The fourth was killed at 07:44:45-07:44:55 together with EVERY run in the repository, 20 of 20 across 7 branches, `windows-msvc-cpu` dying mid-build after passing two steps. That is an Actions-side event, not a verdict on any diff. `gh pr checks` renders a cancelled job as `fail`. My watcher reported 16 failures and the operator's reported 20; the true count both times was ZERO. A per-check listing cannot separate "this branch is red" from "the pool was killed" -- only the run-level `conclusion` can. Carried in the spec because anyone subtracting a CI baseline on this repo will otherwise attribute an infrastructure event to a diff. A re-run was triggered at 07:47 and was still queued at 07:50. Whatever it reports is the CI result; this commit does not claim one. Record-only; no code change. `df -h /` 84% used / 73G free. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [Claude Code]

… against the corrected bound (#496) (#675) FOLLOWING_AGENTS_PROTOCOL W2 of KERNEL-SSM-MAMBA (#496): the CUDA arm of vt::Mamba2ChunkScan, vt::Mamba2StateUpdate and vt::RmsNormGatedGroup. Supersedes #566 and #592, which had diverged -- one carried the evidence resolving the merge precondition, the other the F1/F2 repair, and neither contained the other. F1: the declared equivalence contract named the elementary functions as the ONLY admitted source of device-vs-host divergence. Host C++ is pinned -ffp-contract=off and nothing passes --fmad=false to nvcc, so acc += xv * bv contracts on device and not on host. Stated model 3.5*K*u, contraction adds ~1.0*K*u, bound was 4*(K+2)*u -- provable only for K <= 18 while the driver shapes run at K = 200. Repaired by carrying the term to 5*(K+2)*u with the arithmetic shown in the header, all three test comments and 8.3. -fmad=false was rejected deliberately: it is a per-TU flag on a header included by the hot GDN TU. Repo-wide gap filed as #591. F2: the header claimed the device kernels stay memory-safe under a contract violation, but the validator checks metadata shape and dtype only. Two dropped checks were memory-unsafe -- an out-of-bounds WRITE past a cudaMallocAsync allocation and an out-of-bounds READ of initial_states -- and the stated reason for omitting them (a D2H plus a stream sync) did not apply, since both values are already in registers and the decode kernel does exactly that clamp for free. Repaired with both the clamps and a narrowed claim enumerating all six dropped checks. A third hole the review did not name was found and closed: seq_idx[0] < 0 indexes passed at chunk -1. All four owed gates discharged by measurement, operator-run on the gate host, each stamping its own lock-acquire time, load and disk. nvcc compile for sm_121a clean with CUTLASS, FA2 and Marlin confirmed ENABLED in the configure log rather than assumed. The three CUDA arms 12/2095, 10/5965, 12/3723 all SUCCESS. compute-sanitizer memcheck ERROR SUMMARY 0 errors. And the 9-mutation re-sweep against the MOVED bound: 9 of 9 CAUGHT. The re-sweep is the item a reader would most likely have waved through, and M6 is why it could not be: it aborts at exit 134 while printing "assertions: 2577 | 2577 passed | 0 failed" -- a clean assertions line on a FAILING run -- and is caught only because the harness reads the exit code. Two of the original eight mutations did not COMPILE under -Werror=all-warnings and were being scored as caught; a mutation that will not build is a suite that never ran. test_minimax_h3 is attributed rather than waived. Reproduced standalone on an idle box under the lock at TEST_EXIT=139, it is #486 with root cause #516, signature-for-signature. The independent baseline that PASSED was row/pool-device-key, the branch that FIXES #516, so the baseline carried a fix and this branch does not carry a defect. CI is REMOTE_UNVERIFIED, not green: every run on the predecessor branches ended cancelled, including a repo-wide mass cancellation of 20 runs across 7 branches, and cuda-fat-build never completed -- which is why the compile was run directly on the gate host. Windows reds are the main baseline (#514, #584). The row stays INVENTORIED. No lifecycle move, no measurement claimed, no performance result: a host reference plus its device arm is not a speed number. Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [Claude Code]

localai-bot · 2026-08-13T21:25:57Z

Superseded by #675, which consolidates this branch with the evidence commits from #566 and lands the row.

mudler added 15 commits August 12, 2026 23:00

merge: origin/main into row/KERNEL-SSM-MAMBA-SSD-W2-FINISH (#496)

fe06959

FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [Claude Code]

WIP spec(KERNEL-SSM-MAMBA): re-author $8.3 equivalence contract, clos…

f9915f8

…e $8.2 decode SUBCASE (#496) FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [Claude Code]

WIP spec(KERNEL-SSM-MAMBA): update $8 Now for the W2 state (#496)

08b6e0b

FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [Claude Code]

WIP spec(KERNEL-SSM-MAMBA): record the W2 evidence -- 9/9 mutations c…

f1929de

…aught, sanitizer clean, Debug arm green (#496) FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [Claude Code]

spec(KERNEL-SSM-MAMBA): record the full-ctest result and the one unat…

62a4bad

…tributed failure (#496) FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [Claude Code]

spec(KERNEL-SSM-MAMBA): the ctest ran under undetected GPU contention…

23bc605

… -- a second job locks /tmp/gpu.lock (#496) FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [Claude Code]

spec(KERNEL-SSM-MAMBA): the Windows CI reds are the main baseline; th…

1e81914

…e attribution re-run is REMOTE_UNVERIFIED (#496) FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [Claude Code]

merge: origin/main into the W2 tightening branch

f800e78

FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [Claude Code]

This was referenced Aug 13, 2026

video_engine.cpp reaches Windows with unguarded POSIX stat — reddens windows-msvc-* on EVERY open PR #664

Closed

W2 (KERNEL-SSM-MAMBA): the Mamba2 SSD CUDA arm — 9/9 mutations caught against the corrected bound (#496) #675

Merged

localai-bot closed this Aug 13, 2026

localai-bot mentioned this pull request Aug 13, 2026

fix(#664): the video registry's existence probes stop reaching Windows with POSIX stat #677

Closed

Navigation Menu

fix(KERNEL-SSM-MAMBA): the derived bound omitted nvcc FMA contraction, and the memory-safety claim outran the kernels (#496) - #592

fix(KERNEL-SSM-MAMBA): the derived bound omitted nvcc FMA contraction, and the memory-safety claim outran the kernels (#496)#592
localai-bot wants to merge 15 commits into
mainfrom
row/KERNEL-SSM-MAMBA-SSD-W2-FIX

localai-bot commented Aug 13, 2026 •

edited

Loading

Uh oh!

localai-bot commented Aug 13, 2026

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

2 participants

Conversation

localai-bot commented Aug 13, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

F1 (MEDIUM) — the derived bound's derivation omitted nvcc FMA contraction

F2 (MEDIUM) — the memory-safety claim outran the kernels

F3 / F4 (LOW, record accuracy)

F5 (LOW) — taken

Evidence (implementer, CPU box)

The CUDA arm COMPILES under real nvcc — operator-run, not CI

⚠️ The red checks on this PR are CANCELLED, not FAILED

omitted_gates — blocked by the GPU LOCK, not the host

Uh oh!

localai-bot commented Aug 13, 2026

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

2 participants

localai-bot commented Aug 13, 2026 •

edited

Loading

`omitted_gates` — blocked by the GPU LOCK, not the host